Add PTODSL A5 DSL ST coverage#886
Merged
zhangstevenunity merged 3 commits intoJul 6, 2026
Merged
Conversation
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request introduces native tensor_view and partition_tensor_view folding support in the FoldTileBufIntrinsics pass, updates ExpandTileOp to include view shape and strides in the specialization key, and adds a pto_level parameter to @pto.jit to forward build-level overrides to ptoas. Additionally, VPTOSplitCVModule is updated to normalize sections in-place for pre-annotated modules. Feedback on the changes highlights a concurrency violation in FoldTileBufIntrinsics where a FuncOp pass queries the parent module's symbol table, a limitation in traceViewChain that fails on nested partitions, and an inefficient cleanup loop that should be optimized using a worklist-based dead code elimination approach.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
Codex Review该评论由 review 机器人自动更新。
SummaryReview failed at stage Findings未生成结构化 findings，因为 review 过程提前失败。 Log Tail |
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
9 times, most recently
from
9f6aa25 to
93ec308
Compare
Zhendong404
suggested changes
Jul 3, 2026
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
2 times, most recently
from
4b25c38 to
26da009
Compare
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
from
26da009 to
8aa2570
Compare
Zhendong404
suggested changes
Jul 6, 2026
| PTOAS_BIN="${PTOAS_BIN:-${ROOT_DIR}/build/tools/ptoas/ptoas}" | ||
| PTOAS_FLAGS="${PTOAS_FLAGS:---pto-arch a5 --pto-backend=vpto}" | ||
| # set he HOST_RUNNER to "ssh root@localhost" if must change user to root to access the device | ||
| HOST_RUNNER="${HOST_RUNNER:-}" |
Contributor
There was a problem hiding this comment.
没懂为什么要修改vpto validation相关的2个脚本，这个PR并没有改任何vpto的st用例
|
|
||
| A_TILE_ADDR = 0 | ||
| B_TILE_ADDR = 4096 | ||
| C_TILE_ADDR = 8192 |
Contributor
There was a problem hiding this comment.
PTOAS支持自动分配tile地址，作为打样的话可以引导后续开发者往这个方向写，应该能简化很多东西
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
3 times, most recently
from
d3bf2ce to
e114140
Compare
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
from
e114140 to
445f437
Compare
jimmychou0
force-pushed
the
codex/ptodsl-a5-dsl-st-validation
branch
from
445f437 to
c35bac0
Compare
Zhendong404
approved these changes
Jul 6, 2026
A3 板测完成（有跳过）
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Abstract
This PR adds the first PTODSL-authored A5 DSL ST coverage and updates the PTODSL simulator CI path so the new cases are actually built and run.
The branch was rebuilt after dropping the earlier broad backend workaround. The remaining backend change is intentionally narrow:
FoldTileBufIntrinsicsnow performs fixpoint cleanup of dead view chains exposed after tile intrinsic folding. It deletes only use-empty bridge casts, memref view ops,pto.make_tensor_view/pto.partition_view, and dead tile allocations; it does not rerun fullPTOViewToMemref, does not broadenExpandTileOp, and does not change live view lowering semantics.Problem scenarios covered:
taddvalidates a basictload + tadd + tstorepath outside the oldtilelang_stharness.tload_storevalidates GM view construction,tload,tstore, and layout variants.tcolexpandandtcolsumcover non-trivial tile shapes, valid rows/cols, and tile-op expansion/runtime behavior.tmatmulvalidates a cube tile matmul path, while the existingcube_matrix_pipeline.pyandgemv_mx_pipeline.pyremain part of the simulator suite.mode="explicit"kernels need to compile through PTOAS level3 and should not implicitly enable sync insertion..soartifacts.@pto.simd/@pto.cubeshould not create redundant section wrappers; explicit kind mismatches should fail early with a clear diagnostic.make_tensor_view/partition_viewchains can otherwise leave high-level or memref view ops that later VPTO emission validation rejects.Implementation changes:
test/dsl-st/npu_a5cases fortadd,tload_store,tcolexpand,tcolsum, andtmatmul.predicate_pack.pyandvmulscvt.pyto avoid level3 live memref subviews.kernel_kindwas explicitly authored in@pto.jitwhile preserving the historical default effective kind ofvectorwhen omitted.mode="explicit"toptoas --pto-level=level3; keep explicit mode from implicitly enabling insert-sync.FoldTileBufIntrinsicsand a focused VPTO lit regression.tools/ptoas/driver.cppfor non-debug output paths.torch/torch_npuruntime instead of installing them each run, isolate PTODSL build artifacts, and ensuretest/dsl-st/npu_a5is covered.Validation
Validated on the 144 simulator environment under
/home/zhoujiaming/ptoas-sim-ci/pr886-cleanupusing the LLVM21 VPTO build and CANN simulator:Results:
llvm-lit:764/764passed.ptodsl/tests/test_jit_compile.py: passed.scripts/sim_dsl.sh --soc-version Ascend950PR_9599 test/dsl-st: all cases passed, includingcube_matrix_pipeline,gemv_mx_pipeline,predicate_pack,simt_gm_memory_core,vmulscvt, and the newnpu_a5directory coverage.Local checks:
Also checked that the final diff no longer touches the old broad-workaround files such as
ExpandTileOp.cpp,PTOInstantiateAndInlineOpLib.cpp,Passes.td,tools/ptoas/ptoas.cpp,VPTOOps.td,VPTO.cpp, orVPTOPtrNormalize.cpp.